#include <stdio.h>
#include "run.h"
__global__ void checkIndex(int n_boxes)
{
    int threadsPerBlock=64;
    // 计算列块的数量，用于处理共享内存中的去除标记。
    const int col_blocks = 3;
    // 获取线程在块中的索引。
    const int tid = threadIdx.x;

    // 共享内存数组，用于标记哪些框应该被去除。
    // 注意：数组大小在运行时确定，由CUDA运行时分配。
    extern __shared__ unsigned long long removed[];

    // 初始化去除标记数组。
    for (int i = tid; i < col_blocks; i += blockDim.x) {
        removed[i] = 0;
    }
    // 确保所有线程都完成了数据加载
    __syncthreads();

  // 遍历列块，处理每个块中的线程块。
  for (int nblock = 0; nblock < col_blocks; ++nblock) {
    // 获取当前块的去除值。
    auto removed_val = removed[nblock];
    // 确保所有线程都完成当前块的处理后再继续。
    __syncthreads();
    // 计算当前块的起始偏移量。
    const int i_offset = nblock * threadsPerBlock;
    printf("i_offset:%d\n",i_offset);
    //每个线程执行64次循环
    for (int inblock = 0; inblock < threadsPerBlock; ++inblock) {
      // 计算当前线程在块中的索引。
      const int i = i_offset + inblock;
      printf("i:%d,线程:%d\n",i,tid);
      // 确保索引不超过框的数量。
      if (i >= n_boxes) break;
      // 如果当前线程对应的框未被标记为去除，则进行处理。
      if (!(removed_val & (1ULL << inblock))) {
        // 线程0负责标记该输出框为保留。
        if (tid == 0) {
          //keep[i] = true;
        }
        // 获取当前框的掩码地址。
        
        //auto p = dev_mask + i * col_blocks;
        // 去除所有与当前框重叠的框。
        // 前面的框已经处理过了，只需要比较后面的block里面的框就可以
        for (int j = tid; j < col_blocks; j += blockDim.x) {
            if(j>=nblock){
                printf("i:%d,j:%d,nblock:%d\n",i,j,nblock);
            }
          //if (j >= nblock) removed[j] |= p[j];
          printf("此处循环执行了%d,blocdim.x:%d,线程为:%d\n",j,blockDim.x,tid);
        }
        // 确保所有线程都完成当前块的更新后再继续。
        __syncthreads();
        // 更新当前块的去除值。
        removed_val = removed[nblock];
      }
    }
  }
}

int main(void)
{
    int boxes_num = 192;
    int threadsPerBlock=64;
    int col_blocks = (boxes_num + threadsPerBlock - 1) / threadsPerBlock;
    int col_blocks_alloc = GET_BLOCKS(boxes_num, threadsPerBlock);
    //int col_blocks_alloc = 2;
    printf("%d,%d,%d,%d\n",boxes_num,threadsPerBlock,col_blocks,col_blocks_alloc);
    dim3 block(1);
    dim3 threads(3);

    printf("block: %d, %d, %d\n", block.x, block.y, block.z);
    printf("thread: %d, %d, %d\n", threads.x, threads.y, threads.z);
    
    //printf("grid: %d, %d, %d\n", grid.x, grid.y, grid.z);

    // //checkIndex<<<grid,block>>>();
    checkIndex<<<1,col_blocks,col_blocks * sizeof(unsigned long long)>>>(boxes_num);
    cudaDeviceReset();

    return 0;
}